昨天我們把 LCN 的題目定義清楚了。
給一張 graph，把每個 node 放到平面上，每條 edge 畫成直線。最後不是只看總共有幾個交叉，而是看最慘的那條 edge：
c[e] = edge e 被其他 edges 穿過的次數
K = max(c[e])
從今天開始進到解法。
整個 solver 大致有三個計算問題：
Crossing
給定 layout，算出每條 edge 的 crossing count
Layout
先產生一份比較合理的 node positions
Simulated Annealing
反覆移動 node，搜尋更低的 K
這幾天會照同一個方式拆：先寫最直接、容易驗證的版本，再看 GPU 上真正浪費的是計算、資料搬運，還是同步。
今天先處理 Crossing，而且故意從最慢的版本開始。
因為 Layout 和 SA 最後都要問同一件事：
這份 node positions 到底有多好？
force-directed layout 產生一份新布局，要計分。
SA 搬動一個 node，要比較移動前後的分數。
多條 chain 同時搜尋，每一條也都要計分。
所以 Crossing evaluator 不是最後才跑一次的 validator。它會出現在搜尋的內圈，而且可能被呼叫幾百萬次。
一次 evaluator 慢 10 us
× 一百萬次 proposal
= 10 秒
這就是為什麼第一個 GPU 案例選 Crossing：它的單次計算不華麗，但呼叫頻率很高；每次少做一點，會一路放大到整個 solver。
兩條 edge：
edge e = (a, b)
edge f = (c, d)
其中 a、b、c、d 都是 node id。真正參與幾何運算的是它們的座標：
A = pos[a]
B = pos[b]
C = pos[c]
D = pos[d]
如果兩條 edge 共用 node，就不是我們要計算的 crossing：
if a == c or a == d or b == c or b == d:
return False
剩下使用 Day 21 的方向判斷：
def orient(p, q, r):
return (
(q.x - p.x) * (r.y - p.y)
- (q.y - p.y) * (r.x - p.x)
)
def opposite(x, y):
return (x < 0 and y > 0) or (x > 0 and y < 0)
def proper_crossing(a, b, c, d):
return (
opposite(orient(a, b, c), orient(a, b, d))
and opposite(orient(c, d, a), orient(c, d, b))
)
不用 slope，也不用除法。
如果改用斜率：
slope = (y2 - y1) / (x2 - x1)
垂直線會遇到除以零，浮點誤差也會讓「剛好共線」變得麻煩。方向判斷只用整數加減乘，對整數網格比較合適。
但整數也不是自動安全。
(bx - ax) * (cy - ay)
座標差再相乘，32-bit 很容易溢位。所以 CUDA 版本的中間值使用 long long。這裡的第一個教訓不是 GPU 技巧，而是：幾何 predicate 一旦錯，後面平行化只會更快地產生錯誤答案。
假設 graph 有 E 條 edges。
最簡單的 CPU 版本只掃每個無序 edge pair 一次：
def full_crossing_counts(points, edges):
counts = [0] * len(edges)
for e in range(len(edges)):
for f in range(e):
if proper_crossing(points, edges[e], edges[f]):
counts[e] += 1
counts[f] += 1
return counts
如果 e 和 f 交叉，兩條 edge 的 count 都要加一。
這段的 edge-pair 數量是：
pair_count = E * (E - 1) / 2
所以時間複雜度是：
O(E²)
例如：
E = 100 -> 4,950 pairs
E = 1,000 -> 499,500 pairs
E = 10,000 -> 49,995,000 pairs
edge 數增加十倍，pair 數接近增加一百倍。
這不是微調幾行 C++ 就會消失的成本。
因為它有三個很好的性質。
第一，容易讀。
每一對 edge 都問一次
交叉就讓兩邊各加一
第二，容易建立 invariants。
每一個 crossing 會被記到兩條 edge：
sum(counts) = 2 * crossing_pair_count
所以：
assert sum(counts) % 2 == 0
第三，它不依賴後面任何優化資料結構。
不需要 spatial grid，不需要 incident-edge cache，也不需要保存上一個 layout 的 crossing counts。這讓它適合當 CPU oracle。
後面的 GPU 版本不論怎麼改，都要逐條 edge 對上這份答案：
gpu_counts[e] == cpu_reference_counts[e]
只比較 K 不夠。中間某條 edge 算錯，但剛好沒有改變最大值，測試仍然可能通過。
把剛才的 CPU 程式直接搬到 GPU，最自然的分工是：
thread 0 -> 計算 c[0]
thread 1 -> 計算 c[1]
thread 2 -> 計算 c[2]
...
kernel 大致長這樣：
__global__ void full_counts(
const int* node_xy,
const int* edge_uv,
int num_edges,
int* output
) {
int e = blockIdx.x * blockDim.x + threadIdx.x;
if (e >= num_edges) return;
int count = 0;
for (int f = 0; f < num_edges; ++f) {
count += proper_crossing(node_xy, edge_uv, e, f);
}
output[e] = count;
}
和 CPU 版本相比，它故意把 (e, f) 和 (f, e) 各算一次。
CPU 版本：
只算 (e, f)
同時寫 counts[e] 和 counts[f]
GPU 版本：
thread e 算 (e, f)，只寫 counts[e]
thread f 算 (f, e)，只寫 counts[f]
GPU 多做了接近兩倍的 pair tests，卻換到一個很有價值的性質：
每個 output 只有一個 writer
因此不需要：
atomicAdd(&counts[e], 1);
atomicAdd(&counts[f], 1);
這是一個有點反直覺的 GPU 取捨：有時候多算一次，比讓大量 threads 爭搶同一批 atomic counters 更容易得到穩定吞吐。
但這不是「重複計算一定比 atomic 快」的通則。哪一個比較好，仍然要看 graph 大小、crossing 密度、contention 和實際 profile。今天採用 one-writer 版本，是因為它簡單、deterministic，而且最適合當第一顆可驗證 kernel。
一個 pair test 要讀：
edge e 的兩個 node ids
edge f 的兩個 node ids
四個 nodes 的 x/y
然後做：
shared-endpoint check
bounding-box rejection
four orientation tests
crossing decision
這不是 GEMM 那種規則的 dense arithmetic。
memory access：依 node id 間接讀取座標
control flow：很多 edge pairs 很早就被排除
arithmetic：少量 integer operations
reuse：同一個 e 會重複使用，f 則一路掃過整張圖
因此它很可能先碰到的是資料讀取與分支效率，不是 Tensor Core 算力。
這也是為什麼今天不急著加 shared memory。
如果 graph 很小，把座標合作搬進 shared memory 可能有用；但小 graph 的 kernel launch 和同步成本也很明顯。如果 graph 很大，完整 node/edge arrays 又放不進 shared memory。
在還不知道 reuse pattern 前就把資料全塞進 shared memory，只是把程式寫複雜，不一定真的減少 DRAM traffic。
case_4/examples/crossing_demo.py 會用固定 seed 建立 graph，先跑 CPU oracle，再呼叫 CUDA full_counts：
python examples/crossing_demo.py \
--gpu \
--nodes 48 \
--edges 96 \
--proposals 32 \
--seed 42
這裡的 32 個 proposals 都從同一份 base layout 出發。今天的 full kernel 會對每份 candidate 完整重算。
驗證不是：
CPU K == GPU K
而是：
assert gpu_counts.shape == cpu_counts.shape
assert gpu_counts[e] == cpu_counts[e] for every edge e
實測中，GPU full counts 和 CPU oracle 逐條 edge 一致。
今天先不把 CPU wall time 和 GPU kernel time 硬除成 speedup。CPU 路徑包含 Python oracle、資料準備與 assertion；GPU 紀錄是 warm-up 後的 CUDA event kernel 區間。量測範圍不同，除出來的數字沒有意義。
現在我們有一顆正確而且很好懂的 GPU kernel：
one thread owns one output edge
每條 edge 掃過所有 partner edges
沒有 output atomic contention
結果逐條 edge 對上 CPU oracle
但算法仍然是：
每一個 candidate
-> 重新計算所有 edge pairs
SA 的一次 proposal 通常只移動一個 node。
真正改變位置的，只有接到這個 node 的 edges。其他大部分 edge pairs 根本沒有變。
所以明天不會繼續微調這顆 O(E²) kernel，而是改問題：
既然只搬一個 node，
哪些 crossing 必須重算，
哪些可以直接沿用？
這會把優化從「用更多 GPU threads 做相同工作」，推到「利用 graph 結構減少工作量」。
mur mur時間----
這幾天參加活動，認真講 學習到很多。
我覺得自己很笨，但老實講，能覺得自己很笨是一件很幸福的事。
你應該要覺得很幸福，有這個權利可以覺得自己很笨，
因為這代表這世界還有這麼多美好
可以讓你去學習、去挑戰、去迎接、去面對、去創造。
當然我們還是有很多限制跟瓶頸，有很多不可能每件事都做到的情況，
但你回過頭來想想就會意識到：
Wow, it's purely amazing, like damn! 如果你真心只是想挑戰、學習或找樂趣的話，你會有找不完的樂趣，因為厲害的人真的太多太多了。
這是我這幾天特別大的收穫。
厲害其實有很多層面：
• 厲害不是 AI 用得好就厲害，也不是 AI 用得不好就不厲害
• 不是純手工自己把程式碼打出來就厲害，也不是靠 AI 打出一堆爛 code 就厲害
什麼是厲害？小時候大家應該都有過對厲害的定義，就是發現
「疑 他做得到我做不到的事，幹，超屌超厲害！」
這是我人生第一個 aha moment：哇，他超屌，我好像做不到。
面對這種情況通常有兩種心態：
我覺得這兩種心態都不是我的。
我的心態是：
「哇他很厲害，我想知道他有多厲害，為什麼很厲害，
想知道自己有沒有機會也能做到。」
這不是一定或絕對，更像是我想要一步一步去理解這件事情是怎麼發生的。
為什麼事情發生的當下我站在不理解的角色？
我要怎麼把自己的位置轉換到理解的陣營、甚至變成貢獻者的陣營，
去把問題更好地解決？
這才是我覺得的
厲害。
我講了好多廢話。有人勸我應該要寫 Blog，我有在寫 GPU 系列啊！
至於這種廢文就當作 GPU 文章後面的 murmur 吧。
如果有人支持的話，你也可以只看我的 murmur，不看 Blog、不看 GPU 哈哈。
已閱讀 6 8 7 1 字
iThome鐵人賽
看影片追技術